#include "../inc/ff_utils_ocl.h"

namespace feifei
{
	typedef struct PlatformInfoType
	{
		std::string profile;
		std::string version;
		std::string vendor;
	}T_PlatformInfo;	
	typedef struct DeviceInfoType
	{
		cl_device_type DeviceType;
		cl_uint VendorId;
		std::string Name;

		cl_uint maxItemDim;
		size_t  maxGroupSize;
		size_t  maxItemSize[3];

		cl_uint CuNum;
		cl_uint CoreClkFreqHz;
		cl_uint GlobalMemWidth;
		cl_uint AddrAlign;
		cl_uint minDataTypeAlign;

		size_t maxAllocMemSize;
		cl_ulong GlobalMemSize;
		cl_ulong LocalMemSize;
		size_t CacheMemSize;
		size_t maxParamSize;
		cl_uint maxConstArgs;
		cl_uint maxSamplers;
		size_t timerResolution;

		cl_uint CachelineSize;
		cl_ulong GlobalCacheSize;
		cl_ulong ConstBuffSize;

		cl_bool enEcc;
		cl_bool enUnified;
		cl_bool isLittleEndian;
		cl_bool isAvailable;
		cl_bool isCompilerAvailable;

		cl_bool enImage;
		cl_uint maxReadImageArgs;
		cl_uint maxWriteImageArgs;
		size_t  maxImage2DWidth;
		size_t  maxImage2DHeight;
		size_t  maxImage3DWidth;
		size_t  maxImage3DHeight;
		size_t  maxImage3DDepth;

		std::string DriverVersion;
		std::string OclCVersion;

		cl_device_mem_cache_type MemCacheType;
		cl_device_local_mem_type LocalMemType;
		cl_device_fp_config FpConfig;
		cl_device_exec_capabilities ExecCap;
		cl_command_queue_properties CmdQueueProp;
	}T_DeviceInfo;

    static cl_context Context = 0;
    static cl_command_queue CommandQueue = 0;
	static cl_platform_id PlatformId;
    static cl_uint DeviceCount;
    static int SellectDeviceIndex = 0;
    static cl_device_id DeviceId;
	static T_PlatformInfo PlatformInfo;
	static T_DeviceInfo DeviceInfo;

E_ffState CreateContext()
{
    cl_int errNum;
    cl_uint numPlatforms;
    cl_platform_id firstPlatformId;

    // First, select an OpenCL platform to run on.  For this example, we
    // simply choose the first available platform.  Normally, you would
    // query for all available platforms and select the most appropriate one.
    errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
    if (errNum != CL_SUCCESS || numPlatforms <= 0)
    {
        std::cerr << "Failed to find any OpenCL platforms." << std::endl;
        return E_ffState::RTN_ERR;
    }

    // Next, create an OpenCL context on the platform.  Attempt to
    // create a GPU-based context, and if that fails, try to create
    // a CPU-based context.
    cl_context_properties contextProperties[] =
    {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)firstPlatformId,
        0
    };
    Context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum);
    if (errNum != CL_SUCCESS)
    {
        std::cout << "Could not create GPU context, trying CPU..." << std::endl;
        Context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, NULL, NULL, &errNum);
        if (errNum != CL_SUCCESS)
        {
            std::cerr << "Failed to create an OpenCL GPU or CPU context." << std::endl;
            return E_ffState::RTN_ERR;
        }
    }

    return E_ffState::SUCCESS;
}
E_ffState CreateCommandQueue()
{
    CommandQueue = clCreateCommandQueue(Context, DeviceId, 0, NULL);
    if (CommandQueue == NULL)
    {
        ERR("Failed to create commandQueue for device");
        return E_ffState::RTN_ERR;
    }

    return E_ffState::SUCCESS;
}
void GetPlatformInfo()
{
	cl_int errNum;
	std::size_t sz;
	char* info;

	errNum = clGetPlatformInfo(PlatformId, CL_PLATFORM_PROFILE, 0, NULL, &sz);
	info = (char*)alloca(sizeof(char) * sz);
	errNum = clGetPlatformInfo(PlatformId, CL_PLATFORM_PROFILE, sz, info, NULL);
	PlatformInfo.profile = std::string(info);

	errNum = clGetPlatformInfo(PlatformId, CL_PLATFORM_VERSION, 0, NULL, &sz);
	info = (char*)alloca(sizeof(char) * sz);
	errNum = clGetPlatformInfo(PlatformId, CL_PLATFORM_VERSION, sz, info, NULL);
	PlatformInfo.version = std::string(info);

	errNum = clGetPlatformInfo(PlatformId, CL_PLATFORM_VENDOR, 0, NULL, &sz);
	info = (char*)alloca(sizeof(char) * sz);
	errNum = clGetPlatformInfo(PlatformId, CL_PLATFORM_VENDOR, sz, info, NULL);
	PlatformInfo.vendor = std::string(info);
}
void GetDeviceInfo()
{
	std::size_t sz;
	char* info;

	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(DeviceInfo.CuNum), &DeviceInfo.CuNum, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(DeviceInfo.CoreClkFreqHz), &DeviceInfo.CoreClkFreqHz, NULL);
	DeviceInfo.CoreClkFreqHz = DeviceInfo.CoreClkFreqHz * (cl_uint)1e6;
	clGetDeviceInfo(DeviceId, CL_DEVICE_ADDRESS_BITS, sizeof(DeviceInfo.GlobalMemWidth), &DeviceInfo.GlobalMemWidth, NULL);

	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_PARAMETER_SIZE, sizeof(DeviceInfo.maxParamSize), &DeviceInfo.maxParamSize, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_CONSTANT_ARGS, sizeof(DeviceInfo.maxConstArgs), &DeviceInfo.maxConstArgs, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_SAMPLERS, sizeof(DeviceInfo.maxSamplers), &DeviceInfo.maxSamplers, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(DeviceInfo.maxAllocMemSize), &DeviceInfo.maxAllocMemSize, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(DeviceInfo.ConstBuffSize), &DeviceInfo.ConstBuffSize, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(DeviceInfo.GlobalMemSize), &DeviceInfo.GlobalMemSize, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, sizeof(DeviceInfo.GlobalCacheSize), &DeviceInfo.GlobalCacheSize, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(DeviceInfo.LocalMemSize), &DeviceInfo.LocalMemSize, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, sizeof(DeviceInfo.CachelineSize), &DeviceInfo.CachelineSize, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(DeviceInfo.AddrAlign), &DeviceInfo.AddrAlign, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_ERROR_CORRECTION_SUPPORT, sizeof(DeviceInfo.enEcc), &DeviceInfo.enEcc, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof(DeviceInfo.enUnified), &DeviceInfo.enUnified, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_ENDIAN_LITTLE, sizeof(DeviceInfo.isLittleEndian), &DeviceInfo.isLittleEndian, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, sizeof(DeviceInfo.minDataTypeAlign), &DeviceInfo.minDataTypeAlign, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_AVAILABLE, sizeof(DeviceInfo.isAvailable), &DeviceInfo.isAvailable, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_COMPILER_AVAILABLE, sizeof(DeviceInfo.isCompilerAvailable), &DeviceInfo.isCompilerAvailable, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_PROFILING_TIMER_RESOLUTION, sizeof(DeviceInfo.timerResolution), &DeviceInfo.timerResolution, NULL);

	clGetDeviceInfo(DeviceId, CL_DEVICE_TYPE, sizeof(DeviceInfo.DeviceType), &DeviceInfo.DeviceType, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_VENDOR_ID, sizeof(DeviceInfo.VendorId), &DeviceInfo.VendorId, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(DeviceInfo.FpConfig), &DeviceInfo.FpConfig, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, sizeof(DeviceInfo.MemCacheType), &DeviceInfo.MemCacheType, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_LOCAL_MEM_TYPE, sizeof(DeviceInfo.LocalMemType), &DeviceInfo.LocalMemType, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_EXECUTION_CAPABILITIES, sizeof(DeviceInfo.ExecCap), &DeviceInfo.ExecCap, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_QUEUE_PROPERTIES, sizeof(DeviceInfo.CmdQueueProp), &DeviceInfo.CmdQueueProp, NULL);

	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(DeviceInfo.maxItemDim), &DeviceInfo.maxItemDim, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(DeviceInfo.maxItemSize), &DeviceInfo.maxItemSize, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(DeviceInfo.maxGroupSize), &DeviceInfo.maxGroupSize, NULL);

	clGetDeviceInfo(DeviceId, CL_DEVICE_IMAGE_SUPPORT, sizeof(DeviceInfo.enImage), &DeviceInfo.enImage, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_READ_IMAGE_ARGS, sizeof(DeviceInfo.maxReadImageArgs), &DeviceInfo.maxReadImageArgs, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_MAX_WRITE_IMAGE_ARGS, sizeof(DeviceInfo.maxWriteImageArgs), &DeviceInfo.maxWriteImageArgs, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(DeviceInfo.maxImage2DWidth), &DeviceInfo.maxImage2DWidth, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(DeviceInfo.maxImage2DHeight), &DeviceInfo.maxImage2DHeight, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_IMAGE3D_MAX_WIDTH, sizeof(DeviceInfo.maxImage3DWidth), &DeviceInfo.maxImage3DWidth, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_IMAGE3D_MAX_HEIGHT, sizeof(DeviceInfo.maxImage3DHeight), &DeviceInfo.maxImage3DHeight, NULL);
	clGetDeviceInfo(DeviceId, CL_DEVICE_IMAGE3D_MAX_DEPTH, sizeof(DeviceInfo.maxImage3DDepth), &DeviceInfo.maxImage3DDepth, NULL);

	clGetDeviceInfo(DeviceId, CL_DEVICE_NAME, 0, NULL, &sz);				info = (char*)alloca(sizeof(char) * sz);
	clGetDeviceInfo(DeviceId, CL_DEVICE_NAME, sz, info, NULL);				DeviceInfo.Name = std::string(info);
	clGetDeviceInfo(DeviceId, CL_DRIVER_VERSION, 0, NULL, &sz);				info = (char*)alloca(sizeof(char) * sz);
	clGetDeviceInfo(DeviceId, CL_DRIVER_VERSION, sz, info, NULL);			DeviceInfo.DriverVersion = std::string(info);
	clGetDeviceInfo(DeviceId, CL_DEVICE_OPENCL_C_VERSION, 0, NULL, &sz);	info = (char*)alloca(sizeof(char) * sz);
	clGetDeviceInfo(DeviceId, CL_DEVICE_OPENCL_C_VERSION, sz, info, NULL);	DeviceInfo.OclCVersion = std::string(info);
}
E_ffState InitGpuRuntime()
{
    E_ffState rtn;

    // context
    rtn = CreateContext();
    if (rtn != E_ffState::SUCCESS)
    {
        ERR("Failed to create OpenCL context");
        return E_ffState::RTN_ERR;
    }

	// platform
	clGetPlatformIDs(1, &PlatformId, NULL);
	GetPlatformInfo();

	// device
	clGetDeviceIDs(PlatformId, CL_DEVICE_TYPE_ALL, 0, NULL, &DeviceCount);
	cl_device_id* devices = new cl_device_id[DeviceCount];
	clGetDeviceIDs(PlatformId, CL_DEVICE_TYPE_ALL, DeviceCount,	devices, NULL);
	DeviceId = devices[SellectDeviceIndex];
	delete[] devices;
	GetDeviceInfo();

    // command queue
    rtn = CreateCommandQueue();
    if (rtn != E_ffState::SUCCESS)
    {
        ERR("Failed to create OpenCL command queue");
        return E_ffState::RTN_ERR;
    }
    return E_ffState::SUCCESS;
}
void ReleaseGpuRuntime()
{
    if (CommandQueue != 0)
        clReleaseCommandQueue(CommandQueue);
    if (Context != 0)
        clReleaseContext(Context);
}
void PrintDeviceDetails()
{
	PrintSeperator('-');

	INFO("\t- Device Name: " + DeviceInfo.Name);
	INFO("\t- Driver Version: " + DeviceInfo.DriverVersion);
	INFO("\t-");
	INFO("\t- Multi Processor Number: %d", DeviceInfo.CuNum);
	INFO("\t- Core Clock: " + fmtFreq(DeviceInfo.CoreClkFreqHz));
	INFO("\t-");
	INFO("\t- Global Memory Bus Width: %d(bit)", DeviceInfo.GlobalMemWidth);
	INFO("\t- Global Memory Size: " + fmtSize(DeviceInfo.GlobalMemSize));
	INFO("\t- Local Memory Size: " + fmtSize(DeviceInfo.LocalMemSize));
	INFO("\t-");
	INFO("\t- Max WorkItem Dimensions: %d", DeviceInfo.maxItemDim);
	INFO("\t- Max Group Size: %d", DeviceInfo.maxGroupSize);
	INFO("\t- Max WorkItem Size: (%d, %d, %d)", DeviceInfo.maxItemSize[0], DeviceInfo.maxItemSize[1], DeviceInfo.maxItemSize[2]);

	INFO("\t-");
	INFO("\t- CL_DEVICE_OPENCL_C_VERSION: " + DeviceInfo.OclCVersion);
	INFO("\t- CL_DEVICE_MAX_PARAMETER_SIZE: %d", DeviceInfo.maxParamSize);
	INFO("\t- CL_DEVICE_MAX_CONSTANT_ARGS: %d", DeviceInfo.maxConstArgs);
	INFO("\t- CL_DEVICE_MAX_SAMPLERS: %d", DeviceInfo.maxSamplers);
	INFO("\t-");
	INFO("\t- CL_DEVICE_IMAGE_SUPPORT: %s", DeviceInfo.enImage ? "YES" : "NO");
	INFO("\t- CL_DEVICE_MAX_READ_IMAGE_ARGS: %d", DeviceInfo.maxReadImageArgs);
	INFO("\t- CL_DEVICE_MAX_WRITE_IMAGE_ARGS: %d", DeviceInfo.maxWriteImageArgs);
	INFO("\t- CL_DEVICE_IMAGE2D_MAX_WIDTH: %d", DeviceInfo.maxImage2DWidth);
	INFO("\t- CL_DEVICE_IMAGE2D_MAX_HEIGHT: %d", DeviceInfo.maxImage2DHeight);
	INFO("\t- CL_DEVICE_IMAGE3D_MAX_WIDTH: %d", DeviceInfo.maxImage3DWidth);
	INFO("\t- CL_DEVICE_IMAGE3D_MAX_HEIGHT: %d", DeviceInfo.maxImage3DHeight);
	INFO("\t- CL_DEVICE_IMAGE3D_MAX_DEPTH: %d", DeviceInfo.maxImage3DDepth);
	INFO("\t-");
	INFO("\t- CL_DEVICE_MAX_MEM_ALLOC_SIZE: " + fmtSize(DeviceInfo.maxAllocMemSize));
	INFO("\t- CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE: " + fmtSize(DeviceInfo.ConstBuffSize));
	INFO("\t- CL_DEVICE_GLOBAL_MEM_CACHE_SIZE: " + fmtSize(DeviceInfo.GlobalCacheSize));
	INFO("\t- CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE(Byte): %d", DeviceInfo.CachelineSize);
	INFO("\t- CL_DEVICE_MEM_BASE_ADDR_ALIGN(Byte): %d", DeviceInfo.AddrAlign);
	INFO("\t- CL_DEVICE_PROFILING_TIMER_RESOLUTION: %d", DeviceInfo.timerResolution);
	INFO("\t- CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE(Byte): %d", DeviceInfo.minDataTypeAlign);
	INFO("\t- CL_DEVICE_ERROR_CORRECTION_SUPPORT: %s", DeviceInfo.enEcc ? "YES" : "NO");
	INFO("\t- CL_DEVICE_HOST_UNIFIED_MEMORY: %s", DeviceInfo.enUnified ? "YES" : "NO");
	INFO("\t- CL_DEVICE_ENDIAN_LITTLE: %s", DeviceInfo.isLittleEndian ? "YES" : "NO");
	INFO("\t- CL_DEVICE_COMPILER_AVAILABLE: %s", DeviceInfo.isCompilerAvailable ? "YES" : "NO");
	INFO("\t- CL_DEVICE_AVAILABLE: %s", DeviceInfo.isAvailable ? "YES" : "NO");

	PrintSeperator('-');
}
void PrintDeviceInfoShort()
{
	PrintSeperator('-');

	INFO("\t- Device Name: " + DeviceInfo.Name);
	INFO("\t- Driver Version: " + DeviceInfo.DriverVersion);
	INFO("\t-");
	INFO("\t- Multi Processor Number: %d", DeviceInfo.CuNum);
	INFO("\t- Core Clock: " + fmtFreq(DeviceInfo.CoreClkFreqHz));
	INFO("\t-");
	INFO("\t- Global Memory Bus Width: %d(bit)", DeviceInfo.GlobalMemWidth);
	INFO("\t- Global Memory Size: " + fmtSize(DeviceInfo.GlobalMemSize));
	INFO("\t- Local Memory Size: " + fmtSize(DeviceInfo.LocalMemSize));
	INFO("\t-");
	INFO("\t- Max WorkItem Dimensions: %d", DeviceInfo.maxItemDim);
	INFO("\t- Max Group Size: %d", DeviceInfo.maxGroupSize);
	INFO("\t- Max WorkItem Size: (%d, %d, %d)", DeviceInfo.maxItemSize[0], DeviceInfo.maxItemSize[1], DeviceInfo.maxItemSize[2]);

	PrintSeperator('-');
}
void PrintGpuRuntimeInfo(bool isFullInfo)
{
	INFO("");
	PrintSeperator('=');
	PrintHeader("Runtime", '=');
	PrintSeperator('=');

	// platform
	INFO("- Version: " + PlatformInfo.version);
	INFO("- Vendor: " + PlatformInfo.vendor);

	INFO("- System has %d devices.", DeviceCount);
	INFO("- Sellect device index: %d.", SellectDeviceIndex);

	if (isFullInfo)
	{
		PrintDeviceDetails();
	}
	else
	{
		PrintDeviceInfoShort();
	}
}

void* DevMalloc(size_t byteNum)
{
    cl_mem d_mem = NULL;

    
    d_mem = clCreateBuffer(Context, CL_MEM_READ_WRITE, byteNum, NULL, NULL);
    if (d_mem == NULL)
    {
        std::cerr << "Error creating memory objects." << std::endl;
        return false;
    }

    return (void*)d_mem;
}
void* HstMalloc(size_t byteNum, E_MemType memType)
{
    void* h_mem = malloc(byteNum);
    return h_mem;
}
void DevFree(void * d_mem)
{
    clReleaseMemObject((cl_mem)d_mem);
}
E_ffState MemCopyH2D(void* d_mem, void* h_mem, size_t byteNum)
{
    cl_int errNum;

    errNum = clEnqueueWriteBuffer(CommandQueue, (cl_mem)d_mem, CL_TRUE,
        0, byteNum, h_mem,
        0, NULL, NULL);

    if (errNum != CL_SUCCESS)
    {
        ERR("Error writing result buffer.");
        return E_ffState::RTN_WARN;
    }
    return E_ffState::SUCCESS;
}
E_ffState MemCopyD2H(void* h_mem, void* d_mem, size_t byteNum)
{
    cl_int errNum;

    errNum = clEnqueueReadBuffer(CommandQueue, (cl_mem)d_mem, CL_TRUE,
        0, byteNum, h_mem,
        0, NULL, NULL);

    if (errNum != CL_SUCCESS)
    {
        ERR("Error reading result buffer.");
        return E_ffState::RTN_WARN;
    }
    return E_ffState::SUCCESS;
}
E_ffState MemCopyD2D(void* d_dst, void* d_src, size_t byteNum)
{
    cl_int errNum;

    errNum = clEnqueueCopyBuffer(CommandQueue, (cl_mem)d_src, (cl_mem)d_dst,
        0, 0, byteNum,
        0, NULL, NULL);

    if (errNum != CL_SUCCESS)
    {
        ERR("Error copying result buffer.");
        return E_ffState::RTN_WARN;
    }
    return E_ffState::SUCCESS;
}

enum class BwDir
{
	H2D = 1,
	D2H = 2,
	D2D = 3
};
void bandwidthTest(size_t width, size_t height, E_MemType mem, bool isPitch, BwDir dir, bool isAsync)
{
	size_t pitch = width * sizeof(float);
	size_t dataSize = width * height;
	size_t memSize = dataSize * sizeof(float);
	int iteration = 100;
	float elapsedTimeInMs = 0.0f;
	float bandwidthInMBs = 0.0f;

	float* h_data1, * h_data2;
	cl_mem d_data1, d_data2;

	// 1. allocate device memory
	if (isPitch)
	{
	}
	else
	{
        d_data1 = (cl_mem)DevMalloc(memSize);
        d_data2 = (cl_mem)DevMalloc(memSize);
	}

	// 2. allocate host memory
	switch (mem)
	{
	case E_MemType::Page:
		h_data1 = (float*)malloc(memSize);
		h_data2 = (float*)malloc(memSize);
		break;
	default:
		h_data1 = (float*)malloc(memSize);
		h_data2 = (float*)malloc(memSize);
	}

	// 3. initialize memory
	for (uint64_t i = 0; i < memSize / sizeof(float); i++)
		h_data1[i] = (i & 0xff) * 1.0f;
	if (dir == BwDir::D2H || dir == BwDir::D2D)
		MemCopyH2D(d_data1, h_data1, memSize);

	// 4. copy data
    ffTimer tm;
    tm.Restart();
	for (int i = 0; i < iteration; i++)
	{
		if (isPitch)
		{
		}
		else
		{
			if (isAsync)
			{
			}
			else
			{
				switch (dir)
				{
				case BwDir::H2D:	MemCopyH2D(d_data1, h_data1, memSize); break;
				case BwDir::D2H:	MemCopyD2H(h_data2, d_data1, memSize); break;
				case BwDir::D2D:	MemCopyD2D(d_data2, d_data1, memSize); break;
				}
			}
		}
	}

	// make sure GPU has finished copying
	//cuDeviceSynchronize();

	//get the total elapsed time in ms
    tm.Stop();
    elapsedTimeInMs = (float)tm.ElapsedMilliSec;

	//calculate bandwidth in MB/s
	bandwidthInMBs = ((float)(1 << 10) * memSize * (float)iteration) / (elapsedTimeInMs * (float)(1 << 20));

	if (mem == E_MemType::Page)
	{
		free(h_data1);
		free(h_data2);
	}
	else
	{
	}
    DevFree(d_data1);
    DevFree(d_data2);

	std::string logstr;
	if (isAsync)logstr = "Async";
	else		logstr = "Sync ";
	switch (dir)
	{
	case BwDir::H2D:	logstr += " H -> D"; break;
	case BwDir::D2H:	logstr += " D -> H"; break;
	case BwDir::D2D:	logstr += " D -> D"; break;
	}
	if (isPitch)
	{
		logstr += " Pitch  ";
	}
	else
	{
		logstr += " UnPitch";
	}
	if (dir == BwDir::D2D)
	{
		logstr += "         ";
	}
	else
	{
		switch (mem)
		{
		case E_MemType::Page:	logstr += " pageable"; break;
		}
	}
	LOG("%s %d(MB) = %.3f(GB/s).", logstr.data(), memSize / 1024 / 1024, bandwidthInMBs / 1024);
}
void BandwidthTest()
{
	size_t width = 16 * 1024;
	size_t height = 2 * 1024;
	PrintSeperator('=');
	INFO("BandWidth Test");

	PrintSeperator('-');
	LOG("H->D VS D->H VS D->D");
	bandwidthTest(width, height, E_MemType::Page, false, BwDir::H2D, false);
	bandwidthTest(width, height, E_MemType::Page, false, BwDir::D2H, false);
	bandwidthTest(width, height, E_MemType::Page, false, BwDir::D2D, false);

	PrintSeperator('=');
}


static std::string gKernelPath = "";
std::string get_kernel_path()
{
	if (gKernelPath == "")
		gKernelPath = get_work_path() + "kernel" + DIR_SPT;

	LOG("set kernel folder  : " + gKernelPath);
	return gKernelPath;
}

GpuKernel::GpuKernel(std::string kernelName, std::string srcFile)
{
	this->kernelName = kernelName;
	if (srcFile == "")
		sourceFile = get_kernel_path() + DIR_SPT + kernelName + ".cl";
	else
		sourceFile = srcFile;
	if (creatKernelFromCppFile() != E_ffState::SUCCESS)
		throw(-1);
}
E_ffState GpuKernel::creatKernelFromCppString()
{
	cl_int errNum;

	clProgram = clCreateProgramWithSource(Context, 1,
		(const char**)&sourceString,
		NULL, NULL);
	if (clProgram == NULL)
	{
		ERR("Failed to create CL program from source.");
		return E_ffState::RTN_ERR;
	}

	errNum = clBuildProgram(clProgram, 0, NULL, NULL, NULL, NULL);
	if (errNum != CL_SUCCESS)
	{
		// Determine the reason for the error
		char buildLog[16384];
		clGetProgramBuildInfo(clProgram, DeviceId, CL_PROGRAM_BUILD_LOG,
			sizeof(buildLog), buildLog, NULL);

		std::cerr << "Error in kernel: " << std::endl;
		std::cerr << buildLog;
		clReleaseProgram(clProgram);
		return E_ffState::RTN_ERR;
	}

	// Create OpenCL kernel
	clKernel = clCreateKernel(clProgram, kernelName.c_str(), NULL);
	if (clKernel == NULL)
	{
		ERR("Failed to create kernel.");
		return E_ffState::RTN_ERR;
	}

	return E_ffState::SUCCESS;
}
E_ffState GpuKernel::creatKernelFromCppFile()
{
	std::ifstream kernelFile(sourceFile, std::ios::in);
	if (!kernelFile.is_open())
	{
		ERR("Failed to open file for reading: " + sourceFile);
	}

	std::ostringstream oss;
	oss << kernelFile.rdbuf();
	std::string srcStdStr = oss.str();
	sourceString = (char*)srcStdStr.c_str();

	return creatKernelFromCppString();
}

E_ffState GpuKernel::Launch()
{
	cl_int errNum;
	errNum = clEnqueueNDRangeKernel(CommandQueue, clKernel, 1, NULL,
		GlobalSize.arr(), GroupSize.arr(),
		0, NULL, NULL);
	if (errNum != CL_SUCCESS)
	{
		ERR("Error queuing kernel for execution.");
		return E_ffState::RTN_ERR;
	}

	return E_ffState::SUCCESS;
}

} // end namespace feifei
